#include<iostream>
#include <cuda_runtime.h>
#include <time.h>
#include <chrono>

bool ok(float*target, float*now, int n)
{
    for(int i = 0; i < n; i++)
        if(target[i] != now[i])
            return 0;
    return 1;
}
void init(float * arr, int size){
    time_t t;
    srand((unsigned int) time(&t));

    for(int i = 0; i < size; i++){
        arr[i] = (float)(rand() & 0xFF) / 10.0f;
    }
}

auto sum_common(float * a, float * b, float * c, const int N)
{
    auto start = std::chrono::high_resolution_clock::now();
    for(int i = 0; i < N; i++)
    {
        c[i] = a[i] + b[i];
    }
    auto end = std::chrono::high_resolution_clock::now();
    return end - start;
}


__global__ void vecAdd( const float* a, 
                        const float* b, 
                        float* c, 
                        int n
                      ) //我们可以把__global__后面跟的东西当成线程函数。
{
    /*
    int grid = (N + threads - 1 )/ threads;
    这里的写法是每个float元素都有一个对应的线程，是一种粒度非常小的划分。

    */
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    //blockIdx.x:块号 blockDim.x：块内线程数 threadIdx.x：线程号
    if(idx < n)
    {
        c[idx] = b[idx] + a[idx];
    }
}
float sum_cuda(float * a, float * b, float * c, const int N)
{
    float *c_a,*c_b,*c_c;
    size_t nBs = N * sizeof(float);
    cudaMalloc((float **)&c_a, nBs);
    cudaMalloc((float **)&c_b, nBs);
    cudaMalloc((float **)&c_c, nBs);

    //dst, src, len, methods;
    cudaMemcpy(c_a, a, nBs, cudaMemcpyHostToDevice);
    cudaMemcpy(c_b, b, nBs, cudaMemcpyHostToDevice);



    int threads = 64; //一个block256个线程
    int grid = (N + threads - 1 )/ threads; //启动几个block


    cudaEvent_t t0, t1; //cuda的话用这个几时
    cudaEventCreate(&t0);
    cudaEventCreate(&t1);

    cudaEventRecord(t0);//t0开始计时.当 GPU 执行到流中 cudaEventRecord(t0) 之前的操作时，t0 就会被标记为完成。我觉得这个GPU流可以被理解为一堆排着队等执行的代码

    //做一些要计时的操作
    vecAdd<<<grid, threads>>>(c_a, c_a, c_c, N);

    cudaEventRecord(t1);//告诉 GPU 在执行完所有之前的操作后，标记一个时间点 t1。这个操作本身是异步的，它只是将一个“记录事件”的命令放入 CUDA 流中。CPU 不会等待这个命令真正执行。
    cudaEventSynchronize(t1); // 这是一个同步操作。它会阻塞 CPU，直到 GPU 真正执行到流中 t1 所在的位置，并且 t1 之前的所有操作都已完成。只有当 t1 真正“发生”了，你才能准确地测量它与 t0 之间的时间差。。
    float ms = 0;
    cudaEventElapsedTime(&ms, t0, t1);

    cudaMemcpy(c, c_c, nBs, cudaMemcpyDeviceToHost);

    return ms;

    /*

        cudaMemcpy(c, d_c, nBytes, cudaMemcpyDeviceToHost);

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    cudaEventDestroy(t0); cudaEventDestroy(t1);
    这个是destory的流程
    */

}



__global__ void vecAdd2( const float* a, 
                        const float* b, 
                        float* c, 
                        int n
                      ) //我们可以把__global__后面跟的东西当成线程函数。
{
    // /*这里是粒度较大的分发*/
    int block_id = blockIdx.x;
    int threads_num = blockDim.x; 
    int thread_id = threadIdx.x;//注意啊这个是block里面的线程id，而不是全局的线程id
    int block_num = gridDim.x;
    //int num_per_thread = n / (block_num * threads_num); //这玩意最好在主机一次性算完。
    // int begin_idx = block_id * threads_num + thread_id;
    // int end_idx = begin_idx + num_per_thread;
    // for(int i = begin_idx; i < end_idx; i++)
    // {
    //     c[i] = a[i] + b[i];
    // } 
    //这个有点小问题啊，这个不是常见手段。

    //接下来这个叫步进循环
    int stride = block_num * threads_num;
    for(int i = block_id * threads_num + thread_id; i < n; i += stride){
        c[i] = a[i] + b[i];
    }

}

float sum_cuda2(float * a, float * b, float * c, const int N)
{
    float *c_a,*c_b,*c_c;
    size_t nBs = N * sizeof(float);
    cudaMalloc((float **)&c_a, nBs);
    cudaMalloc((float **)&c_b, nBs);
    cudaMalloc((float **)&c_c, nBs);

    //dst, src, len, methods;
    cudaMemcpy(c_a, a, nBs, cudaMemcpyHostToDevice);
    cudaMemcpy(c_b, b, nBs, cudaMemcpyHostToDevice);



    int threads = 64; //一个block256个线程
    int grid = 4; //启动几个block


    cudaEvent_t t0, t1; //cuda的话用这个几时
    cudaEventCreate(&t0);
    cudaEventCreate(&t1);

    cudaEventRecord(t0);//t0开始计时.当 GPU 执行到流中 cudaEventRecord(t0) 之前的操作时，t0 就会被标记为完成。我觉得这个GPU流可以被理解为一堆排着队等执行的代码

    //做一些要计时的操作
    vecAdd2<<<grid, threads>>>(c_a, c_a, c_c, N);

    cudaEventRecord(t1);//告诉 GPU 在执行完所有之前的操作后，标记一个时间点 t1。这个操作本身是异步的，它只是将一个“记录事件”的命令放入 CUDA 流中。CPU 不会等待这个命令真正执行。
    cudaEventSynchronize(t1); // 这是一个同步操作。它会阻塞 CPU，直到 GPU 真正执行到流中 t1 所在的位置，并且 t1 之前的所有操作都已完成。只有当 t1 真正“发生”了，你才能准确地测量它与 t0 之间的时间差。。
    float ms = 0;
    cudaEventElapsedTime(&ms, t0, t1);

    cudaMemcpy(c, c_c, nBs, cudaMemcpyDeviceToHost);

    return ms;

    /*

        cudaMemcpy(c, d_c, nBytes, cudaMemcpyDeviceToHost);

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    cudaEventDestroy(t0); cudaEventDestroy(t1);
    这个是destory的流程
    */

}


int main(){
    int n = (1 << 10);
    std::cout<<"共" << n << "个\n";

    float *a, *b, *c, *c2, *c3;
    a = new float[n];
    b = new float[n];
    c = new float[n];
    c2 = new float[n];
    c3 = new float[n];

    init(a, n);
    init(b, n);


    std::chrono::duration<double> duration = sum_common(a, b, c, n);
    std::cout << "代码执行耗时: " << duration.count() << " 秒" << std::endl;


    std::cout << "代码执行耗时: " << sum_cuda(a, b, c2, n)  << " 秒" << std::endl;
    std::cout << "是否一样 "<< ok(c, c2, n) << std::endl;

    
    std::cout << "代码执行耗时: " << sum_cuda2(a, b, c3, n)  << " 秒" << std::endl;
    std::cout << "是否一样 "<< ok(c, c3, n) << std::endl;


    delete a;
    delete b;
    delete c;

    return 0;


}





/*
我好像有点理解多维度的分配方式了。比如二维的，可以天然的更合适的应用于图片之类的。


比如我们有个1024 * 1024 * 1024的体数据。我们一个线程块中，每个维度可以有8个线程。如果每个线程对以一个数据点，那么一个线程块可以有
(8,8,8)个数据点。
于是需要(128, 128, 128)个线程块。

于是<<<gridSize(128, 128, 128), blockSize(8, 8, 8)>>>

然后blockIdx的xyz就是访问(128, 128, 128)个线程块中的某个线程块
然后threadIdx的xyz就是访问当前块中(8, 8, 8)个线程中的某个线程

grimDim的xyz是告诉你对应的(128, 128, 128)的信息。这里grimDim是128.
blockDim也是类似的。

 */